#include <iostream>
#include <cstdlib>
#include <fstream>
#include <string>
#include <vector>
#include "cl/opencl.h"
#include "dft.h"

using namespace std;

#define ClCheck(expr) { \
  cl_int code = expr; \
  if (code != CL_SUCCESS) { \
    cout << "OpenCL error: " << #expr << ", code " << code \
         << ", line " << __LINE__ << ", file " << __FILE__ << endl; \
    exit(EXIT_FAILURE); \
  } \
}

namespace openclimpl {

  class OpenClInit {
  public:
    static OpenClInit& Instance() {
      static OpenClInit init;
      return init;
    }

    ~OpenClInit() {
      clReleaseProgram(program);
      clReleaseContext(context);
      clReleaseCommandQueue(queue);
      clReleaseKernel(reorder_kernel);
      clReleaseKernel(dft_kernel);
    }

    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    const char* src;
    cl_kernel reorder_kernel;
    cl_kernel dft_kernel;

  private:
    OpenClInit();

    string str_;
  };


  OpenClInit::OpenClInit() {
    cl_int status;
    cl_device_type dev_type;

    cl_uint num_platforms;
    ClCheck(clGetPlatformIDs(0, 0, &num_platforms));

    vector<cl_platform_id> platform_ids(num_platforms);
    ClCheck(clGetPlatformIDs(num_platforms, &platform_ids[0], 0));

    char name[1024];
    for (cl_uint i = 0; i < num_platforms; ++i) {
      ClCheck(clGetPlatformInfo(platform_ids[i], CL_PLATFORM_NAME, 1024, &name, 0));
      if (strstr(name, "NVIDIA")) {
        cout << "OpenClInit: NVIDIA OpenCL was detected\n";
        dev_type = CL_DEVICE_TYPE_GPU;
        platform = platform_ids[i];
        break;
      }
      //if (strstr(name, "Intel")) {
      //  cout << "OpenClInit: INTEL OpenCL was detected\n";
      //  dev_type = CL_DEVICE_TYPE_CPU;
      //  platform = platform_ids[i];
      //  break;
      //}
    }

    ClCheck(clGetDeviceIDs(platform, dev_type, 1, &device, NULL));

    context = clCreateContext(0, 1, &device, NULL, NULL, &status); 
    ClCheck(status);

    queue = clCreateCommandQueue(context, device, 0, &status);
    ClCheck(status);

    ifstream file("dft.cl");    
    str_ = string((istreambuf_iterator<char>(file)), istreambuf_iterator<char>());
    src = str_.c_str();
    file.close();

    program = clCreateProgramWithSource(context, 1, &src, NULL, &status);
    ClCheck(status);

    ClCheck(clBuildProgram(program, 0, NULL, NULL, NULL, NULL));

    reorder_kernel = clCreateKernel(program, "ReorderElements", &status);
    ClCheck(status);

    dft_kernel = clCreateKernel(program, "DftOnePass", &status);
    ClCheck(status);
  }


  void DftInplace(std::complex<float>* x, int n) {
    OpenClInit& cl = OpenClInit::Instance();

    cl_mem x_d = clCreateBuffer(cl.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, 2*n*sizeof(float), x, NULL);

    size_t global_work_sz = n;
    size_t local_work_sz = min(n, 512);

    int num_bits = int(log(double(n)) / log(2.));
    ClCheck(clSetKernelArg(cl.reorder_kernel, 0, sizeof(x_d), &x_d));
    ClCheck(clSetKernelArg(cl.reorder_kernel, 1, sizeof(int), &num_bits));
    ClCheck(clEnqueueNDRangeKernel(cl.queue, cl.reorder_kernel, 1, NULL, &global_work_sz, &local_work_sz, 0, NULL, NULL));

    global_work_sz = n / 2;
    local_work_sz = min(n / 2, 512);

    ClCheck(clSetKernelArg(cl.dft_kernel, 0, sizeof(x_d), &x_d));
    for (int size = 1, size_log = 0; size < n; size <<= 1, size_log++) {
      ClCheck(clSetKernelArg(cl.dft_kernel, 1, sizeof(size_log), &size_log));
      ClCheck(clEnqueueNDRangeKernel(cl.queue, cl.dft_kernel, 1, NULL, &global_work_sz, &local_work_sz, 0, NULL, NULL));
    }

    ClCheck(clEnqueueReadBuffer(cl.queue, x_d, CL_TRUE, 0, 2*n*sizeof(float), x, 0, NULL, NULL));
    clReleaseMemObject(x_d);
  }

}